/*
 * Projekt: GMU 2013 - Canny hranovy detektor OpenCL/CUDA
 * Autori: Pavel Nevrala - xnevra00 (xnevra00@stud.fit.vutbr.cz)
 *         Martin Jurak - xjurak02 (xjurak02@stud.fit.vutbr.cz)
 */
#include "SobelFilterCL.h"


SobelFilterCL::SobelFilterCL(cl_uint maxWorkGroupP, cl_context *contextP, cl_command_queue *commandQueueP, cl_program *programP): 
	CannyPartsBase(maxWorkGroupP), context(contextP), commandQueue(commandQueueP), program(programP), edgesDirection(NULL)
{	
}


SobelFilterCL::~SobelFilterCL(void)
{
	if(edgesDirection)
		delete[](edgesDirection);
}

/*
 * Provede nastaveni OpenCl a kernelu pro vypocet gradientu obrazu a jejich smeru
 * Sobeluv operator ve smeru X a Y je proveden separabilnimi filtry pro kazdy smer 
 * SobelXX - {1,0,-1} (Sobel oprerator ve smeru X, maska pro radkovy filtr)
 * SobelXY - {1,2,1} (Sobel oprerator ve smeru X, maska pro sloupcovy filtr)
 * SobelYX - {1,2,1} (Sobel oprerator ve smeru Y, maska pro radkovy filtr)
 * SobelYY - {1,0,-1} (Sobel oprerator ve smeru Y, maska pro sloupcovy filtr)
 * pro sloupcove a radkuve filtry je provedena konvoluce s obrazem.
 *
 * @image - vstupni obraz
 * @width - sirka obrazu
 * @height - vyska obrazu	 
 * @return - doba vypoctu na GPU
 */
double SobelFilterCL::run(cl_float *image, cl_uint width, cl_uint height)
{
	cl_int clErr;
	cl_uint maskSize = 3;

	// masky pro gradient ve semru X
	cl_float sobelXX[3] = {1,0,-1};
	cl_float sobelXY[3] = {1,2,1};

	// masky pro gradient ve semru Y
	cl_float sobelYX[3] = {1,2,1};
	cl_float sobelYY[3] = {1,0,-1};

	cl_event event_sobelXX, event_sobelXY,event_sobelYX, event_sobelYY, event_sobelResult;

	//vytvoreni bufferu
	cl_mem cl_in = clCreateBuffer(*context, CL_MEM_READ_ONLY, width * height * sizeof(cl_float), NULL, &clErr);
	clErr = clEnqueueWriteBuffer(*commandQueue, cl_in, CL_TRUE /*blocking write*/,	0, width * height * sizeof(cl_float), image, 0, NULL, NULL);

	cl_mem cl_outXX = clCreateBuffer(*context, CL_MEM_READ_WRITE, width * height * sizeof(cl_float), NULL, &clErr);
	cl_mem cl_outGX = clCreateBuffer(*context, CL_MEM_READ_WRITE, width * height * sizeof(cl_float), NULL, &clErr);
	cl_mem cl_outYX = clCreateBuffer(*context, CL_MEM_READ_WRITE, width * height * sizeof(cl_float), NULL, &clErr);	
	cl_mem cl_outGY = clCreateBuffer(*context, CL_MEM_READ_WRITE, width * height * sizeof(cl_float), NULL, &clErr);

	cl_mem cl_edgesDirection = clCreateBuffer(*context, CL_MEM_WRITE_ONLY, width * height * sizeof(cl_float), NULL, &clErr);
	cl_mem cl_out = clCreateBuffer(*context, CL_MEM_WRITE_ONLY, width * height * sizeof(cl_float), NULL, &clErr);
	cl_mem cl_maskX = clCreateBuffer(*context, CL_MEM_READ_ONLY, 3 * sizeof(cl_float), NULL, &clErr);
	cl_mem cl_maskY = clCreateBuffer(*context, CL_MEM_READ_ONLY, 3 * sizeof(cl_float), NULL, &clErr);

	//ziskani kernelu	
	cl_kernel kernelSobelResult = clCreateKernel(*program, "sobelResult", &clErr);
	errorController.CheckOpenCLError( clErr, "clCreateKernel edge_result" );

	cl_kernel kernelSobelXX = clCreateKernel(*program, "convolutionX", &clErr);
	errorController.CheckOpenCLError( clErr, "clCreateKernel edge_XX" );
	cl_kernel kernelSobelXY = clCreateKernel(*program, "convolutionY", &clErr);
	errorController.CheckOpenCLError( clErr, "clCreateKernel edge_XY" );
	cl_kernel kernelSobelYX = clCreateKernel(*program, "convolutionX", &clErr);
	errorController.CheckOpenCLError( clErr, "clCreateKernel edge_YX" );
	cl_kernel kernelSobelYY = clCreateKernel(*program, "convolutionY", &clErr);
	errorController.CheckOpenCLError( clErr, "clCreateKernel edge_YY" );

//========================================================================================================================================================================================
// ========================================================================================================================================================================================
	
	// maska pro SobelXX
	clErr = clEnqueueWriteBuffer(*commandQueue, cl_maskX, CL_TRUE /*blocking write*/,	0, 3 * sizeof(cl_float), sobelXX, 0, NULL, NULL);
		
	clErr = clSetKernelArg(kernelSobelXX, 0, sizeof(cl_mem), &cl_in);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (inputImage)");
	clErr = clSetKernelArg(kernelSobelXX, 1, sizeof(cl_mem), &cl_outXX);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (outImage)");
	clErr = clSetKernelArg(kernelSobelXX, 2, sizeof(cl_uint), &width);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (widthImage)");
	clErr = clSetKernelArg(kernelSobelXX, 3, sizeof(cl_uint), &height);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (heightImage)");
	clErr = clSetKernelArg(kernelSobelXX, 4, sizeof(cl_float), &cl_maskX);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (mask)");
	clErr = clSetKernelArg(kernelSobelXX, 5, sizeof(cl_uint), &maskSize);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (maskSize)");

	size_t GWS[2], LWS[2];
	GWS[0]=maxWorkGroup;
	GWS[1]=height;
	LWS[0]=maxWorkGroup;
	LWS[1]=1;

	// provedeni konvoluce s radkovou maskou
	clErr = clEnqueueNDRangeKernel(*commandQueue, kernelSobelXX, 2, NULL, GWS, LWS, 0, NULL, &event_sobelXX);
	errorController.CheckOpenCLError(clErr, "enquekernel sobelXX");

	clErr = clWaitForEvents(1, &event_sobelXX);
	errorController.CheckOpenCLError(clErr, "waitforevents sobelX");

// ========================================================================================================================================================================================
// ========================================================================================================================================================================================
	
	// maska pro SobelXY
	cl_event event_sobelMaskXY;
	clErr = clEnqueueWriteBuffer(*commandQueue, cl_maskX, CL_TRUE /*blocking write*/,	0, 3 * sizeof(cl_float), sobelXY, 1,&event_sobelXX, &event_sobelMaskXY);
	
   	clErr = clSetKernelArg(kernelSobelXY, 0, sizeof(cl_mem), &cl_outXX);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (inputImage)");
	clErr = clSetKernelArg(kernelSobelXY, 1, sizeof(cl_mem), &cl_outGX);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (outImage)");
	clErr = clSetKernelArg(kernelSobelXY, 2, sizeof(cl_uint), &width);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (widthImage)");
	clErr = clSetKernelArg(kernelSobelXY, 3, sizeof(cl_uint), &height);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (heightImage)");
	clErr = clSetKernelArg(kernelSobelXY, 4, sizeof(cl_float), &cl_maskX);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (mask)");
	clErr = clSetKernelArg(kernelSobelXY, 5, sizeof(cl_uint), &maskSize);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (maskSize)");

	GWS[0]=width;
	GWS[1]=maxWorkGroup;
	LWS[0]=1;
	LWS[1]=maxWorkGroup;

	cl_event waitingEvents[] = {event_sobelXX, event_sobelMaskXY};

	// provedeni konvoluce s sloupcovou maskou
	clErr = clEnqueueNDRangeKernel(*commandQueue, kernelSobelXY, 2, NULL, GWS, LWS, 2, waitingEvents, &event_sobelXY);//&event_gaussX
	errorController.CheckOpenCLError(clErr, "enquekernelsobelXY");

	clErr = clWaitForEvents(1, &event_sobelXY);
	errorController.CheckOpenCLError(clErr, "waitforevents sobelX");

// ========================================================================================================================================================================================
// ========================================================================================================================================================================================
	
	// maska pro SobelYX
	clErr = clEnqueueWriteBuffer(*commandQueue, cl_maskY, CL_TRUE /*blocking write*/,	0, 3 * sizeof(cl_float), sobelYX, 0,NULL, NULL);	
	
	clErr = clSetKernelArg(kernelSobelYX, 0, sizeof(cl_mem), &cl_in);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (inputImage)");
	clErr = clSetKernelArg(kernelSobelYX, 1, sizeof(cl_mem), &cl_outYX);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (outImage)");
	clErr = clSetKernelArg(kernelSobelYX, 2, sizeof(cl_uint), &width);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (widthImage)");
	clErr = clSetKernelArg(kernelSobelYX, 3, sizeof(cl_uint), &height);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (heightImage)");
	clErr = clSetKernelArg(kernelSobelYX, 4, sizeof(cl_float), &cl_maskY);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (mask)");
	clErr = clSetKernelArg(kernelSobelYX, 5, sizeof(cl_uint), &maskSize);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (maskSize)");
	
	GWS[0]=maxWorkGroup;
	GWS[1]=height;
	LWS[0]=maxWorkGroup;
	LWS[1]=1;

	// provedeni konvoluce s radkovou maskou
	clErr = clEnqueueNDRangeKernel(*commandQueue, kernelSobelYX, 2, NULL, GWS, LWS, NULL, NULL, &event_sobelYX);//&event_gaussX
	errorController.CheckOpenCLError(clErr, "enquekernel sobelYX");

	clErr = clWaitForEvents(1, &event_sobelYX);
	errorController.CheckOpenCLError(clErr, "waitforevents sobelX");

// ========================================================================================================================================================================================
// ========================================================================================================================================================================================
	
	// maska pro SobelYY
	cl_event event_sobelMaskYY;
	clErr = clEnqueueWriteBuffer(*commandQueue, cl_maskY, CL_TRUE /*blocking write*/,	0, 3 * sizeof(cl_float), sobelYY, 1, &event_sobelYX, &event_sobelMaskYY);

   	clErr = clSetKernelArg(kernelSobelYY, 0, sizeof(cl_mem), &cl_outYX);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (inputImage)");
	clErr = clSetKernelArg(kernelSobelYY, 1, sizeof(cl_mem), &cl_outGY);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (outImage)");
	clErr = clSetKernelArg(kernelSobelYY, 2, sizeof(cl_uint), &width);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (widthImage)");
	clErr = clSetKernelArg(kernelSobelYY, 3, sizeof(cl_uint), &height);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (heightImage)");
	clErr = clSetKernelArg(kernelSobelYY, 4, sizeof(cl_float), &cl_maskY);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (mask)");
	clErr = clSetKernelArg(kernelSobelYY, 5, sizeof(cl_uint), &maskSize);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (maskSize)");

	GWS[0]=width;
	GWS[1]=maxWorkGroup;
	LWS[0]=1;
	LWS[1]=maxWorkGroup;

	waitingEvents[0]= event_sobelYX; 
	waitingEvents[1]=	event_sobelMaskYY;

	// provedeni konvoluce s sloupcovou maskou
	clErr = clEnqueueNDRangeKernel(*commandQueue, kernelSobelYY, 2, NULL, GWS, LWS, 2, waitingEvents, &event_sobelYY);//&event_gaussX
	errorController.CheckOpenCLError(clErr, "enquekernelsobelYY");

	clErr = clWaitForEvents(1, &event_sobelYY);
	errorController.CheckOpenCLError(clErr, "waitforevents sobelX");


// ========================================================================================================================================================================================
// ========================================================================================================================================================================================
	
	// vysledna velikost a smer gradientu
   	clErr = clSetKernelArg(kernelSobelResult, 0, sizeof(cl_mem), &cl_outGX);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (inputImageSobelX)");
	clErr = clSetKernelArg(kernelSobelResult, 1, sizeof(cl_mem), &cl_outGY);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (outImageSobelY)");
	clErr = clSetKernelArg(kernelSobelResult, 2, sizeof(cl_mem), &cl_out);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (outImageSobelResult)");
	clErr = clSetKernelArg(kernelSobelResult, 3, sizeof(cl_mem), &cl_edgesDirection);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (outImageEdgesDirection)");
	clErr = clSetKernelArg(kernelSobelResult, 4, sizeof(cl_uint), &width);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (widthImage)");
	clErr = clSetKernelArg(kernelSobelResult, 5, sizeof(cl_uint), &height);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (heightImage)");		

	GWS[0]=maxWorkGroup;
	GWS[1]=height;
	LWS[0]=maxWorkGroup;
	LWS[1]=1;
	
	waitingEvents[0] = event_sobelXY;
	waitingEvents[1] = event_sobelYY;

	// vysledna velikost a smer gradientu
	clErr = clEnqueueNDRangeKernel(*commandQueue, kernelSobelResult, 2, NULL, GWS, LWS, 2, waitingEvents, &event_sobelResult);
	errorController.CheckOpenCLError(clErr, "enquekernelSobelResult");
		
	clErr = clWaitForEvents(1, &event_sobelResult);	
	errorController.CheckOpenCLError(clErr, "waitforevents Sobel");

// =======================================================================================================================================================================================
// ========================================================================================================================================================================================
	
	// zjisteni doby vypoctu
	clGetEventProfilingInfo(event_sobelXX, CL_PROFILING_COMMAND_START, sizeof(time_start), &time_start, NULL);
	clGetEventProfilingInfo(event_sobelResult, CL_PROFILING_COMMAND_END, sizeof(time_end), &time_end, NULL);
	total_time = (double)(time_end - time_start);

	// zapis zpracovaneho obrazu
	outImage = new cl_float[width * height];
	memset(outImage, 0, width * height * sizeof(cl_float));
	clErr = clEnqueueReadBuffer (*commandQueue, cl_out, CL_TRUE, 0, width * height * sizeof(cl_float), outImage, 0, NULL, NULL);
	errorController.CheckOpenCLError(clErr, "enqueread buffer Sobel");

	// zapis obrazu se smery gradientu
	edgesDirection = new cl_float[width * height];
	memset(edgesDirection, 0, width * height * sizeof(cl_float));
	clErr = clEnqueueReadBuffer (*commandQueue, cl_edgesDirection, CL_TRUE, 0, width * height * sizeof(cl_float), edgesDirection, 0, NULL, NULL);
	errorController.CheckOpenCLError(clErr, "enqueread buffer Sobel edges direction");

	// uvolneni zdroju
    clErr = clReleaseKernel(kernelSobelXX);
    errorController.CheckOpenCLError(clErr, "clReleaseKernel edgeXX.");
    clErr = clReleaseKernel(kernelSobelXY);
    errorController.CheckOpenCLError(clErr, "clReleaseKernel edgeXY.");
	clErr = clReleaseKernel(kernelSobelYX);
    errorController.CheckOpenCLError(clErr, "clReleaseKernel edgeYX.");
	clErr = clReleaseKernel(kernelSobelYY);
    errorController.CheckOpenCLError(clErr, "clReleaseKernel edgeYY.");	
	clErr = clReleaseKernel(kernelSobelResult);
    errorController.CheckOpenCLError(clErr, "clReleaseKernel edgeResult.");

	clErr = clReleaseMemObject(cl_in);
    errorController.CheckOpenCLError(clErr, "clReleaseMemObject input"); 
    clErr = clReleaseMemObject(cl_outXX);
    errorController.CheckOpenCLError(clErr, "clReleaseMemObject sobelXX");
    clErr = clReleaseMemObject(cl_outYX);
    errorController.CheckOpenCLError(clErr, "clReleaseMemObject sobelYX");	
	 clErr = clReleaseMemObject(cl_outGX);
    errorController.CheckOpenCLError(clErr, "clReleaseMemObject GX");	
	 clErr = clReleaseMemObject(cl_outGY);
    errorController.CheckOpenCLError(clErr, "clReleaseMemObject GY");	
    clErr = clReleaseMemObject(cl_out);
    errorController.CheckOpenCLError(clErr, "clReleaseMemObject out");
	clErr = clReleaseMemObject(cl_edgesDirection);
    errorController.CheckOpenCLError(clErr, "clReleaseMemObject edgeDirection");

	// doba vypoctu v milisekundach
	return total_time / 1000000.0;

}